/*
 * Printf Debugging Integration for GPU Correctness Verification
 * Provides comprehensive debugging information for thread indices and memory access
 */

#include <cuda_runtime.h>
#include <stdio.h>
#include <stdarg.h>

// Debug level enumeration
enum DebugLevel {
    DEBUG_NONE = 0,
    DEBUG_BASIC = 1,
    DEBUG_DETAILED = 2,
    DEBUG_VERBOSE = 3
};

// Debug configuration structure
struct DebugConfig {
    DebugLevel level = DEBUG_BASIC;
    bool enable_thread_info = true;
    bool enable_memory_access = true;
    bool enable_timing = true;
    int max_prints = 100;  // Limit printf output
    int print_stride = 1000;  // Print every Nth element
};

// Thread debugging structure
struct ThreadDebugInfo {
    int blockIdx_x, blockIdx_y, blockIdx_z;
    int threadIdx_x, threadIdx_y, threadIdx_z;
    int global_thread_id;
    int warp_id;
    int lane_id;
    int sm_id;
    
    __device__ void populate() {
        blockIdx_x = blockIdx.x;
        blockIdx_y = blockIdx.y;
        blockIdx_z = blockIdx.z;
        threadIdx_x = threadIdx.x;
        threadIdx_y = threadIdx.y;
        threadIdx_z = threadIdx.z;
        
        global_thread_id = (blockIdx.x * gridDim.y * gridDim.z + 
                           blockIdx.y * gridDim.z + blockIdx.z) *
                           (blockDim.x * blockDim.y * blockDim.z) +
                           (threadIdx.x * blockDim.y * blockDim.z +
                           threadIdx.y * blockDim.z + threadIdx.z);
        
        warp_id = threadIdx.x / 32;
        lane_id = threadIdx.x % 32;
        
        // Get SM ID (requires inline PTX)
        asm("mov.u32 %0, %%smid;" : "=r"(sm_id));
    }
};

// Memory access debugging structure
struct MemoryDebugInfo {
    void* base_address;
    size_t offset;
    size_t total_size;
    int access_type;  // 0: read, 1: write
    float value;
    
    __device__ void populate(void* base, size_t off, size_t size, int type, float val) {
        base_address = base;
        offset = off;
        total_size = size;
        access_type = type;
        value = val;
    }
};

// Enhanced validation kernel with printf debugging
__global__ void debug_validation_kernel(
    const float* __restrict__ a,
    const float* __restrict__ b,
    float* __restrict__ c,
    int n,
    int* __restrict__ debug_info,
    float* __restrict__ debug_values,
    int debug_offset,
    int debug_level,
    int max_prints,
    int print_stride
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        ThreadDebugInfo thread_info;
        thread_info.populate();
        
        // Debug printing based on level
        if (debug_level >= DEBUG_BASIC && (idx % print_stride == 0)) {
            printf("[DEBUG] Thread[%d,%d,%d][%d,%d,%d] idx=%d\n",
                   thread_info.blockIdx_x, thread_info.blockIdx_y, thread_info.blockIdx_z,
                   thread_info.threadIdx_x, thread_info.threadIdx_y, thread_info.threadIdx_z,
                   idx);
        }
        
        if (debug_level >= DEBUG_DETAILED && (idx < max_prints)) {
            printf("[DETAIL] GlobalID=%d Warp=%d Lane=%d SM=%d\n",
                   thread_info.global_thread_id,
                   thread_info.warp_id,
                   thread_info.lane_id,
                   thread_info.sm_id);
        }
        
        // Perform computation
        float result = a[idx] + b[idx];
        c[idx] = result;
        
        if (debug_level >= DEBUG_VERBOSE && (idx < max_prints)) {
            printf("[VERBOSE] a[%d]=%.6f b[%d]=%.6f result=%.6f\n",
                   idx, a[idx], idx, b[idx], result);
        }
        
        // Store debug information
        if (debug_info && idx < debug_offset) {
            debug_info[idx * 6] = thread_info.blockIdx_x;
            debug_info[idx * 6 + 1] = thread_info.threadIdx_x;
            debug_info[idx * 6 + 2] = thread_info.global_thread_id;
            debug_info[idx * 6 + 3] = thread_info.warp_id;
            debug_info[idx * 6 + 4] = thread_info.lane_id;
            debug_info[idx * 6 + 5] = thread_info.sm_id;
            
            if (debug_values && idx < debug_offset) {
                debug_values[idx * 5] = a[idx];
                debug_values[idx * 5 + 1] = b[idx];
                debug_values[idx * 5 + 2] = result;
                debug_values[idx * 5 + 3] = a[idx] + b[idx]; // Expected
                debug_values[idx * 5 + 4] = (float)idx;      // Original index
            }
        }
    }
}

// Debug masquerading kernel with detailed logging
__global__ void debug_masquerading_kernel(
    const float** __restrict__ a_arrays,
    const float** __restrict__ b_arrays,
    float** __restrict__ c_arrays,
    int array_size,
    int num_sub_tasks,
    int* __restrict__ debug_info,
    float* __restrict__ debug_values,
    int debug_offset,
    int debug_level,
    int max_prints,
    int print_stride
) {
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total_elements = array_size * num_sub_tasks;
    
    if (global_idx < total_elements) {
        ThreadDebugInfo thread_info;
        thread_info.populate();
        
        int sub_task_id = global_idx / array_size;
        int local_idx = global_idx % array_size;
        
        if (debug_level >= DEBUG_BASIC && (global_idx % print_stride == 0)) {
            printf("[DEBUG] Masquerading: Global=%d SubTask=%d Local=%d\n",
                   global_idx, sub_task_id, local_idx);
        }
        
        if (debug_level >= DEBUG_DETAILED && (global_idx < max_prints)) {
            printf("[DETAIL] Thread[%d,%d] processing sub-task %d\n",
                   thread_info.blockIdx_x, thread_info.threadIdx_x, sub_task_id);
        }
        
        if (sub_task_id < num_sub_tasks && local_idx < array_size) {
            const float* a = a_arrays[sub_task_id];
            const float* b = b_arrays[sub_task_id];
            float* c = c_arrays[sub_task_id];
            
            float result = a[local_idx] + b[local_idx];
            c[local_idx] = result;
            
            if (debug_level >= DEBUG_VERBOSE && (global_idx < max_prints)) {
                printf("[VERBOSE] ST=%d L=%d: a=%.6f b=%.6f r=%.6f\n",
                       sub_task_id, local_idx, a[local_idx], b[local_idx], result);
            }
            
            // Store debug information
            if (debug_info && global_idx < debug_offset) {
                debug_info[global_idx * 7] = thread_info.blockIdx_x;
                debug_info[global_idx * 7 + 1] = thread_info.threadIdx_x;
                debug_info[global_idx * 7 + 2] = sub_task_id;
                debug_info[global_idx * 7 + 3] = local_idx;
                debug_info[global_idx * 7 + 4] = thread_info.global_thread_id;
                debug_info[global_idx * 7 + 5] = thread_info.warp_id;
                debug_info[global_idx * 7 + 6] = thread_info.lane_id;
                
                if (debug_values && global_idx < debug_offset) {
                    debug_values[global_idx * 6] = a[local_idx];
                    debug_values[global_idx * 6 + 1] = b[local_idx];
                    debug_values[global_idx * 6 + 2] = result;
                    debug_values[global_idx * 6 + 3] = a[local_idx] + b[local_idx];
                    debug_values[global_idx * 6 + 4] = (float)sub_task_id;
                    debug_values[global_idx * 6 + 5] = (float)local_idx;
                }
            }
        }
    }
}

// Memory access debugging kernel
__global__ void memory_debug_kernel(
    const float* __restrict__ a,
    const float* __restrict__ b,
    float* __restrict__ c,
    int n,
    int* __restrict__ memory_info,
    float* __restrict__ memory_values,
    int debug_offset,
    int debug_level
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        // Memory address calculation
        size_t a_addr = (size_t)(&a[idx]);
        size_t b_addr = (size_t)(&b[idx]);
        size_t c_addr = (size_t)(&c[idx]);
        
        if (debug_level >= DEBUG_DETAILED && (idx < 50)) {
            printf("[MEMORY] idx=%d a@%p b@%p c@%p\n",
                   idx, (void*)a_addr, (void*)b_addr, (void*)c_addr);
        }
        
        // Check memory alignment
        bool a_aligned = (a_addr % 4 == 0);
        bool b_aligned = (b_addr % 4 == 0);
        bool c_aligned = (c_addr % 4 == 0);
        
        if (!a_aligned || !b_aligned || !c_aligned) {
            printf("[WARNING] Unaligned access at idx=%d\n", idx);
        }
        
        float result = a[idx] + b[idx];
        c[idx] = result;
        
        // Store memory information
        if (memory_info && idx < debug_offset) {
            memory_info[idx * 6] = (int)(a_addr & 0xFFFFFFFF);
            memory_info[idx * 6 + 1] = (int)(b_addr & 0xFFFFFFFF);
            memory_info[idx * 6 + 2] = (int)(c_addr & 0xFFFFFFFF);
            memory_info[idx * 6 + 3] = a_aligned ? 1 : 0;
            memory_info[idx * 6 + 4] = b_aligned ? 1 : 0;
            memory_info[idx * 6 + 5] = c_aligned ? 1 : 0;
            
            if (memory_values && idx < debug_offset) {
                memory_values[idx * 4] = a[idx];
                memory_values[idx * 4 + 1] = b[idx];
                memory_values[idx * 4 + 2] = result;
                memory_values[idx * 4 + 3] = (float)((a_addr - (size_t)a) / sizeof(float));
            }
        }
    }
}

// Debug configuration class
class DebugManager {
private:
    DebugConfig config;
    
public:
    DebugManager(const DebugConfig& cfg = DebugConfig()) : config(cfg) {}
    
    void set_debug_level(DebugLevel level) { config.level = level; }
    void set_max_prints(int max_prints) { config.max_prints = max_prints; }
    void set_print_stride(int stride) { config.print_stride = stride; }
    
    DebugLevel get_debug_level() const { return config.level; }
    int get_max_prints() const { return config.max_prints; }
    int get_print_stride() const { return config.print_stride; }
    
    // Launch debug validation kernel
    void launch_debug_validation(
        const float* a, const float* b, float* c, int n,
        int* debug_info, float* debug_values, int debug_offset
    ) {
        int threads_per_block = 256;
        int blocks = (n + threads_per_block - 1) / threads_per_block;
        
        debug_validation_kernel<<<blocks, threads_per_block>>>(
            a, b, c, n, debug_info, debug_values, debug_offset,
            config.level, config.max_prints, config.print_stride
        );
    }
    
    // Launch debug masquerading kernel
    void launch_debug_masquerading(
        const float** a_arrays, const float** b_arrays, float** c_arrays,
        int array_size, int num_sub_tasks,
        int* debug_info, float* debug_values, int debug_offset
    ) {
        int total_elements = array_size * num_sub_tasks;
        int threads_per_block = 256;
        int blocks = (total_elements + threads_per_block - 1) / threads_per_block;
        
        debug_masquerading_kernel<<<blocks, threads_per_block>>>(
            a_arrays, b_arrays, c_arrays, array_size, num_sub_tasks,
            debug_info, debug_values, debug_offset,
            config.level, config.max_prints, config.print_stride
        );
    }
    
    // Launch memory debug kernel
    void launch_memory_debug(
        const float* a, const float* b, float* c, int n,
        int* memory_info, float* memory_values, int debug_offset
    ) {
        int threads_per_block = 256;
        int blocks = (n + threads_per_block - 1) / threads_per_block;
        
        memory_debug_kernel<<<blocks, threads_per_block>>>(
            a, b, c, n, memory_info, memory_values, debug_offset, config.level
        );
    }
    
    // Print debug summary
    void print_debug_summary() {
        printf("Debug Configuration:\n");
        printf("  Level: %d\n", config.level);
        printf("  Max prints: %d\n", config.max_prints);
        printf("  Print stride: %d\n", config.print_stride);
    }
};

// C-style interface functions
extern "C" {
    DebugManager* create_debug_manager(int debug_level, int max_prints, int print_stride) {
        DebugConfig config;
        config.level = static_cast<DebugLevel>(debug_level);
        config.max_prints = max_prints;
        config.print_stride = print_stride;
        return new DebugManager(config);
    }
    
    void destroy_debug_manager(DebugManager* manager) {
        delete manager;
    }
    
    void set_debug_level(DebugManager* manager, int level) {
        manager->set_debug_level(static_cast<DebugLevel>(level));
    }
    
    void set_debug_max_prints(DebugManager* manager, int max_prints) {
        manager->set_max_prints(max_prints);
    }
    
    void set_debug_print_stride(DebugManager* manager, int stride) {
        manager->set_print_stride(stride);
    }
}

// Utility function for debugging memory
__device__ void debug_memory_access(const char* prefix, int idx, 
                                   const float* ptr, float value) {
    size_t addr = (size_t)ptr;
    printf("[%s] idx=%d addr=%p value=%.6f\n", prefix, idx, (void*)addr, value);
}